feat(cuda-device): add SharedPtr, a CTA-shared pointer typed addrspace(3) - #1341

…l TMA without SharedPtr (#1656)

## What This PR Does

This PR makes three changes to shared-memory addressing and CTA-local
TMA loads (following up on #1336, #1341, #1238, #1337, and #1249):

1. **Keep `.add()` and `.cast()` on shared pointers in `addrspace(3)`
(`mir-importer`)**
2. **Use 32-bit `i32` shared-memory offsets in `cuda-oxide-mlir-export`
(+ `#[grid_constant]` in `blockscale_gemm_cute`)**
3. **Remove `SharedPtr<T>` (`-3,141` lines) and take `dst: *mut u8`
directly in `cp_async_bulk_tensor_{1..5}d_g2s_cta` (+ route `ctm-rs` to
`.shared::cta`)**

---

### 1. Keep `.add()` and `.cast()` on Shared Pointers in `addrspace(3)`
(`mir-importer`)

Previously, `SlotAddrSpaceMap` only preserved `addrspace(3)` across
plain `_y = _x` copies. Calling `.add(offset)` or `.cast::<T>()` on a
shared pointer demoted it to a 64-bit generic pointer (`addrspace(0)`),
forcing an `addrspace(3) -> addrspace(0) -> addrspace(3)` round-trip:

```text
Rust kernel:
  let base = DynamicSharedArray::<u8, 1024>::get_raw();
  let tile = base.add(slot * TILE_BYTES).cast::<u32>();

BEFORE:
  base (ptr addrspace(3))
    ──► addrspacecast to generic (ptr)       <-- demoted on .add() / .cast()!
    ──► 64-bit generic GEP
    ──► addrspacecast back to shared (ptr addrspace(3))

AFTER:
  base (ptr addrspace(3))
    ──► GEP in ptr addrspace(3)              <-- stays in shared memory (0 casts)!
```

---

### 2. Use 32-bit `i32` Shared-Memory Offsets in
`cuda-oxide-mlir-export`

Shared memory (`!llvm.ptr<3>`) uses 32-bit addresses, but
`MirPtrOffsetOp` and SM120 `mma` / `epilogue` lowering in
`cuda-oxide-mlir-export` were extending 32-bit lane/warp indices to
64-bit `i64` before indexing `!llvm.ptr<3>`.

```text
BEFORE (64-bit i64 index math on 32-bit !llvm.ptr<3>):
  lane (i32) ──► zext to i64 ──► 64-bit swizzle math ──► GEP(!llvm.ptr<3>, i64)
  * Blinded LLVM's LoadStoreVectorizer (emitted 4 scalar 32-bit LDS instructions)
  * Left 20 64-bit SHF.R.U64 / IADD.64 instructions in the K-loop

AFTER (32-bit i32 index math on 32-bit !llvm.ptr<3>):
  lane (i32) ──────────────────► 32-bit swizzle math ──► GEP(!llvm.ptr<3>, i32)
  * LLVM's LoadStoreVectorizer automatically merges the 4 scalar loads into LDS.128!
  * Eliminates the 64-bit shift/add instructions in SASS
```

---

### 3. Remove `SharedPtr<T>` and Take `dst: *mut u8` in `.shared::cta`
TMA (#1336, #1341, #1238, #1337, #1249)

#### Why #1238 Happened on RTX 5090 (`sm_120a`)
Consumer Blackwell (`sm_120a`) has no hardware support for cross-block
TMA loads. When a kernel emits a **Cluster TMA** load
(`.shared::cluster`), `ptxas` generates an `if / else` with a software
emulation `CALL` in the `else` branch, which allocates an 8-byte stack
frame and causes `ptxas` to ignore `setmaxnreg` (warning `C7506`):

```text
Cluster TMA (`.shared::cluster`) on sm_120a:
  if (ptr.block_id == my_block_id) {
      UTMALDG.2D ;                                                 // fast hardware TMA
  } else {
      CALL __cuda_syscall_cp_async_bulk_tensor_2d_tile_unicast ;   // <-- kills setmaxnreg!
  }

Local Block TMA (`.shared::cta`) on sm_120a:
  UTMALDG.2D ;   // no branch, no CALL, 0 stack frame, keeps setmaxnreg!
```

#### Why We Don't Need `SharedPtr<T>` (Revising #1336 / #1341)
PR #1336 (and #1341) introduced `SharedPtr<T>` and lowered it to `ptr
addrspace(3)` across function signatures, return values, structs, and
arrays (`+3,141` lines of plumbing). That created a fundamental layout
mismatch:
* `rustc` always sizes `SharedPtr<T>` at **8 bytes** (`*mut T`).
* Under `p3:32:32:32` (`libNVVM` / `--nvptx-short-ptr`), LLVM sizes `ptr
addrspace(3)` at **4 bytes**, causing return-type mismatches in `llc`
and `misaligned address` crashes in arrays/structs.

With Change 1 (`SlotAddrSpaceMap`) + `convert_g2s_cta` (which already
accepts both `addrspace(3)` and `addrspace(0)` pointers, just like
`barrier: *mut Barrier`), `cp_async_bulk_tensor_{1..5}d_g2s_cta` works
directly with `dst: *mut u8`:

```rust
// Direct CTA-local TMA load with plain *mut u8 (0 cvta.to.shared, 0 stack frame):
cp_async_bulk_tensor_2d_g2s_cta(
    &raw mut TILE as *mut u8,
    tensor_map,
    tile_x,
    tile_y,
    &raw mut BAR,
);
```

We also updated `ctm_rs::tma::cp_async_bulk_tensor_shared_cta_global` to
call `cp_async_bulk_tensor_2d_g2s_cta` (`.shared::cta`) instead of
`cp_async_bulk_tensor_2d_g2s` (`.shared::cluster`), eliminating the
8-byte stack frame and runtime CTA-id check in `blockscale_gemm_ctm`.

---

## Verification

* **All 238 examples pass `scripts/smoketest.sh` on RTX 5090
(`sm_120a`)**, including `tma_copy/verify-code-shape.sh` (`.shared::cta`
emitted, `0` `cvta.to.shared`, `0` `shared::cluster`).
* **`just check`** (`fmt-check`, `clippy`, `test`, `test-cuda`,
`check-guards`, `check-intrinsics main`, `doc-check`): **all PASS**.
* **`scripts/check-primitives-api-compile.py`** (`api`, `fp16`,
`blockscale`): **all PASS**.
* **`setmaxnreg` + `.shared::cta` on `sm_120a`** (`ptxas -v` +
`nvdisasm`): `168` registers, `0` stack frame,
`USETMAXREG.DEALLOC.CTAPOOL 0x28`, `USETMAXREG.TRY_ALLOC.CTAPOOL UP0,
0xe8`, `UTMALDG.2D`, no `CALL`.
* **2048x2048x1024 MXFP4 Block-Scaled GEMM on RTX 5090 (Nsight `nsys
profile` kernel-active median of 31 launches)**:

| Kernel | Backend | Before This PR | **After This PR** | vs. Python
CUTLASS CTM (`10.880 us`) |

|:-------|:--------|:---------------|:------------------|:-------------------------------------|
| **Python CUTLASS CTM** (`reference/benchmark.py`) | `cutlass-mlir` |
`10.880 us` | **`10.880 us`** | baseline |
| **Rust `blockscale_gemm_cute`** | `cutlass-mlir` | `11.360 us` |
**`10.912 us`** | **`-0.032 us` (`-0.3%`)** |
| **Rust `blockscale_gemm_ctm`** | **Native `LLVM/PTX`** | `11.616 us` |
**`10.976 us`** | **`-0.096 us` (`-0.9%`)** |

---------

Signed-off-by: nihalpasham <[email protected]>